#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

using namespace std;

static __global__ void testKernel(int *array, int* sumOutput, int n) {
    printf("array is : %s\n", __isLocal(array) ? "local" : (__isGlobal(array) ? "global" : "none"));  // global
    int pos = blockIdx.x * blockDim.x + threadIdx.x;
    extern __shared__ int cache[];  // 动态shareMemory，kernelLaunch时传入shareMemory大小（单位byte）

    int value = 0;
    int numThreadPerBlock = blockDim.x;
    int currThreadIdx = threadIdx.x;

    if (pos < n) {
        value = array[pos];
    }

    for (int i = numThreadPerBlock / 2; i > 0; i /= 2) {
        cache[currThreadIdx] = value;
        // __syncthreads();  // 同步等待block中其他线程
        if (currThreadIdx <= i) value += cache[currThreadIdx + i];
        // __syncthreads();  //  
    }

    if (currThreadIdx == 0) {
        printf("block %d value = %d\n", blockIdx.x, value);
        atomicAdd(sumOutput, value);  // sumOutput是global，value是local
    }

}


// 给定具有n个元素的数组array，对所有元素求和并打印
int main() {
    const int num = 10;
    int* array = nullptr;
    int* sumOutput = nullptr;
    // __host__​cudaError_t cudaMallocManaged ( void** devPtr, size_t size, unsigned int  flags = cudaMemAttachGlobal )
    // 1.**统一内存空间**：
    // - 单个指针同时在CPU和GPU上有效
    // - 无需显式的cudaMemcpy操作
    // 2.**按需迁移**：
    // - 数据在CPU或GPU首次访问时自动迁移
    // - 支持页面级粒度迁移 (从Pascal架构开始)
    // 3.**一致性模型**：
    // - 默认弱一致性 (需要同步点)
    // - 可使用cudaDeviceSynchronize()强制同步
    cudaMallocManaged(&array, sizeof(int) * num);
    cudaMallocManaged(&sumOutput, sizeof(int));
    *sumOutput = 0;

    int groundTruth = 0;
    for (int i = 0; i < num; ++i) {
        array[i] = i;
        groundTruth += i;
    }

    // 请让numThreadPerBlock是2的幂次方，因为block的规约是2倍步长缩减
    int numThreadPerBlock = 8;  // 每个block 8个线程
    int numBlockPerGrid = (num + numThreadPerBlock - 1) / numThreadPerBlock;
    printf("numBlockPerGrid = %d, numThreadPerBlock = %d\n", numBlockPerGrid, numThreadPerBlock);  // num个元素需要多少个block

    int sharedBytes = sizeof(int) * numThreadPerBlock;
    testKernel<<<numBlockPerGrid, numThreadPerBlock, sharedBytes>>>(array, sumOutput, num);

    cudaDeviceSynchronize();
    printf("sumOutput = %d, groundTruth = %d\n", *sumOutput, groundTruth);

    return 0;
}